#include "filter_gpu_global.h"
#include <assert.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>
#include <device_functions.h> //_synchronizedThreads()
#include <stdlib.h>
#include <stdio.h>


/**
 *
 * @param img image
 * @param width size X
 * @param sizeX size window X
 * @param sizeY size window Y
 * @return media of the pixels in the window
 */
__device__ uchar4 promedio(uchar4* img, int width, int sizeX, int sizeY){
	int X =blockIdx.x*DIMBLOCK+threadIdx.x;
	int Y = blockIdx.y*DIMBLOCK+threadIdx.y;
	int RGB[3]={0};
	uchar4 pixel;
	for(int j=Y-sizeY/2;j<=Y+sizeY/2;j++){
		for(int i=X-sizeX/2;i<=X+sizeX/2;i++){
			RGB[0]+=img[width*j+i].x;
			RGB[1]+=img[width*j+i].y;
			RGB[2]+=img[width*j+i].z;
		}
	}
	pixel.x=RGB[0]/(sizeX*sizeY);
	pixel.y=RGB[1]/(sizeX*sizeY);
	pixel.z=RGB[2]/(sizeX*sizeY);
	return pixel;
}
/**
 *
 * @param buf filter image
 * @param img image
 * @param width size in X
 * @param height size in Y
 * @param x position in X
 * @param y position in Y
 */
__global__ void mediaFilter_global(uchar4* buf, uchar4* img, int width, int height, int x,int y){
	int X =blockIdx.x*DIMBLOCK+threadIdx.x;
	int Y = blockIdx.y*DIMBLOCK+threadIdx.y;
	if(X>=x/2 && X<width-x/2 && Y<height-y/2 && Y>=y/2){
		buf[Y*width+X]= promedio(img,width,x,y);
	}
	else{
		buf[Y*width+X].x= 0;
		buf[Y*width+X].y= 0;
		buf[Y*width+X].z= 0;
	}
}

void mediaFilter_gpu(uchar4* h_img,int width, int height, int kernelSizeX, int kernelSizeY){

	assert(h_img!=NULL && width!=0 && height!=0);

	//DEVICE
	uchar4* d_img;
	uchar4* d_buf;
	//ALOCAR MEMORIA HOST

	//ALOCAR MEMORIA DEVICE

	checkCudaErrors(cudaMalloc((void**) &(d_img), width*height*sizeof(unsigned char)*4));
	checkCudaErrors(cudaMalloc((void**) &(d_buf), width*height*sizeof(unsigned char)*4));

	//COPY IMG TO DEVICE

	checkCudaErrors(cudaMemcpy(d_img, h_img, width*height*sizeof(uchar4), cudaMemcpyHostToDevice));
	dim3 dimGrid((width+DIMBLOCK-1) /DIMBLOCK,(height+DIMBLOCK-1) / DIMBLOCK,1);
	dim3 threads(DIMBLOCK,DIMBLOCK,1);

	mediaFilter_global<<<dimGrid,threads>>>(d_buf,d_img,width,height, kernelSizeX,kernelSizeY);

	checkCudaErrors(cudaMemcpy(h_img, d_buf, width*height*sizeof(uchar4), cudaMemcpyDeviceToHost));

	checkCudaErrors(cudaFree(d_img));
	checkCudaErrors(cudaFree(d_buf));

}
